﻿// OCLGlobalAvgPool.cpp : 此文件包含 "main" 函数。程序执行将在此处开始并结束。
//

#include <iostream>
#include "OpenCLKernel.h"
#include <math.h>
#include <time.h>


const char* _kernel_global_avg_pool = KERNEL(
	__kernel void avg_pool(int imgSize, global float* imgsIn, global float* imgsOut) {

	const int li0 = get_local_id(0);

	int offset = li0 * imgSize;

	float sum = 0.0;
	for (int i = 0; i < imgSize; i++) {
		sum += imgsIn[offset + i];
	}
	imgsOut[li0] = sum / imgSize;
}
);


bool floatEq(const float a, const float b) {
	return fabs(a - b) < 0.0001;
}

bool check(int depth, int imgSize, float* inImg, float* outImg) {

	float sum = 0.0;
	for (int i = 0; i < imgSize; i++) {
		sum += inImg[i];
	}
	return floatEq(outImg[depth], sum / imgSize);

}


int globalAvgPool()
{
	int imgSize = 7 * 7;
	static const int depth = 256;

	float* inImg[depth];
	for (int i = 0; i < depth; i++) {
		inImg[i] = (float*)malloc(sizeof(float) * imgSize);
		srand(2);
		for (int j = 0; j < imgSize; inImg[i][j] = rand() & 0xF, j++);
	}
	std::cout << depth << " inImgs is generated!\n";

	float* outImg = (float*)malloc(sizeof(float) * depth);

	cl_int status;

	OpenCLKernel* m_pOCL = new OpenCLKernel;
	if (0 != m_pOCL->OpenCLInitialize())
	{
		delete m_pOCL;
		m_pOCL = nullptr;
		return EXIT_FAILURE;
	}

	//拼接
	float* allInImgs = (float*)malloc(sizeof(float) * imgSize * depth);

	for (int d = 0; d < depth; d++) {
		for (int j = 0; j < imgSize; j++) {
			allInImgs[d * imgSize + j] = inImg[d][j];
		}
	}

	cl_command_queue queue = clCreateCommandQueue(m_pOCL->GetContext(), m_pOCL->GetDevices()[0], NULL, &status);
	GET_CL_ERROR(status, m_pOCL);

	//编译
	size_t nSourceSize = NULL;
	cl_program program = NULL;
	cl_kernel kernel = NULL;


	nSourceSize = { strlen(_kernel_global_avg_pool) };
	program = clCreateProgramWithSource(m_pOCL->GetContext(), 1, &_kernel_global_avg_pool, &nSourceSize, &status);
	GET_CL_ERROR(status, m_pOCL);
	status = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
	GET_CL_ERROR(status, m_pOCL);
	kernel = clCreateKernel(program, "avg_pool", &status);
	GET_CL_ERROR(status, m_pOCL);

	//读内存
	cl_mem bufferIn = clCreateBuffer(m_pOCL->GetContext(), CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * imgSize * depth, allInImgs, &status);
	GET_CL_ERROR(status, m_pOCL);
	cl_mem bufferOut = clCreateBuffer(m_pOCL->GetContext(), CL_MEM_WRITE_ONLY, sizeof(float) *  depth, NULL, &status);
	GET_CL_ERROR(status, m_pOCL);

	clock_t time_start = clock();


	//参数赋值
	status = clSetKernelArg(kernel, 0, sizeof(int), &imgSize);
	GET_CL_ERROR(status, m_pOCL);
	status = clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufferIn);
	GET_CL_ERROR(status, m_pOCL);
	status = clSetKernelArg(kernel, 2, sizeof(cl_mem), &bufferOut);
	GET_CL_ERROR(status, m_pOCL);


	// 注意不同函数需要调整工作组网格参数
	size_t* globalSize = NULL;
	size_t* localSize = NULL;
	cl_uint work_dim = 0;

	work_dim = 1;
	globalSize = new size_t[1];
	globalSize[0] = depth;


	//执行
	status = clEnqueueNDRangeKernel(queue, kernel, work_dim, NULL, globalSize, localSize, 0, NULL, NULL);
	GET_CL_ERROR(status, m_pOCL);
	status = clFinish(queue);
	GET_CL_ERROR(status, m_pOCL);

	clock_t time_end = clock();
	printf("\nTime kernel : %d ms\n", time_end - time_start);

	//写内存
	status = clEnqueueReadBuffer(queue, bufferOut, CL_TRUE, 0, sizeof(float) * depth, outImg, 0, NULL, NULL);
	GET_CL_ERROR(status, m_pOCL);


	// 返回并检查结果
	printf("Checking...\n");
	bool correct = 1;
	for (int i = 0; correct && i < depth; i++) {
		correct = check(i, imgSize, inImg[i], outImg);
		if (!correct) {
			printf("error occur at img %d\n", i);
		}
	}

	if (correct)
		printf("Result correct.\n");

	//结束并释放
	if (globalSize != NULL)
		delete[] globalSize;
	if (localSize != NULL)
		delete[] localSize;
	clReleaseMemObject(bufferIn);
	clReleaseMemObject(bufferOut);
	clReleaseKernel(kernel);
	clReleaseProgram(program);
	clReleaseCommandQueue(queue);
	delete m_pOCL;
	for (int i = 0; i < depth; i++)
		free(inImg[i]);
	free(outImg);

	return 0;
}





int main()
{
	globalAvgPool();
}

// 运行程序: Ctrl + F5 或调试 >“开始执行(不调试)”菜单
// 调试程序: F5 或调试 >“开始调试”菜单

// 入门使用技巧: 
//   1. 使用解决方案资源管理器窗口添加/管理文件
//   2. 使用团队资源管理器窗口连接到源代码管理
//   3. 使用输出窗口查看生成输出和其他消息
//   4. 使用错误列表窗口查看错误
//   5. 转到“项目”>“添加新项”以创建新的代码文件，或转到“项目”>“添加现有项”以将现有代码文件添加到项目
//   6. 将来，若要再次打开此项目，请转到“文件”>“打开”>“项目”并选择 .sln 文件
